





<!DOCTYPE html>
<html class="writer-html5" lang="zh-CN" >
<head>
  <meta charset="utf-8">
  
  <meta name="viewport" content="width=device-width, initial-scale=1.0">
  
  <title>如何在GPU上优化卷积 &mdash; tvm 0.8.dev1982 文档</title>
  

  
  <link rel="stylesheet" href="https://maxcdn.bootstrapcdn.com/bootstrap/4.0.0/css/bootstrap.min.css" integrity="sha384-Gn5384xqQ1aoWXA+058RXPxPg6fy4IWvTNh0E263XmFcJlSAwiGgFAW/dAiS6JXm" crossorigin="anonymous">
  <link rel="stylesheet" href="../../_static/css/theme.css" type="text/css" />
  <link rel="stylesheet" href="../../_static/pygments.css" type="text/css" />
  <link rel="stylesheet" href="../../_static/css/theme.css" type="text/css" />
  <link rel="stylesheet" href="../../_static/gallery.css" type="text/css" />
  <link rel="stylesheet" href="../../_static/pygments.css" type="text/css" />
  <link rel="stylesheet" href="../../_static/css/tlcpack_theme.css" type="text/css" />

  
  
    <link rel="shortcut icon" href="../../_static/tvm-logo-square.png"/>
  

  
  
  
  
    
      <script type="text/javascript" id="documentation_options" data-url_root="../../" src="../../_static/documentation_options.js"></script>
        <script data-url_root="../../" id="documentation_options" src="../../_static/documentation_options.js"></script>
        <script src="../../_static/jquery.js"></script>
        <script src="../../_static/underscore.js"></script>
        <script src="../../_static/doctools.js"></script>
        <script src="../../_static/translations.js"></script>
    
    <script type="text/javascript" src="../../_static/js/theme.js"></script>

    
    <script type="text/javascript" src="../../_static/js/tlcpack_theme.js"></script>
    <link rel="index" title="索引" href="../../genindex.html" />
    <link rel="search" title="搜索" href="../../search.html" />
    <link rel="next" title="How to optimize convolution using TensorCores" href="opt_conv_tensorcore.html" />
    <link rel="prev" title="如何在CPU上优化GEMM（通用矩阵乘）" href="opt_gemm.html" /> 
</head>

<body class="wy-body-for-nav">

   
  <div class="wy-grid-for-nav">
    
    
<header class="header">
    <div class="innercontainer">
      <div class="headerInner d-flex justify-content-between align-items-center">
          <div class="headerLogo">
               <a href="https://tvm.apache.org/"><img src=https://tvm.apache.org/assets/images/logo.svg alt="logo"></a>
          </div>

          <div id="headMenu" class="headerNav">
            <button type="button" id="closeHeadMenu" class="navCloseBtn"><img src="../../_static/img/close-icon.svg" alt="Close"></button>
             <ul class="nav">
                <li class="nav-item">
                   <a class="nav-link" href=https://tvm.apache.org/community>Community</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://tvm.apache.org/download>Download</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://tvm.apache.org/vta>VTA</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://tvm.apache.org/blog>Blog</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://tvm.apache.org/docs>Docs</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://tvmconf.org>Conference</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://github.com/apache/tvm/>Github</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://tvmchinese.github.io/declaration_zh_CN.html>About-Translators</a>
                </li>
             </ul>
               <div class="responsivetlcdropdown">
                 <button type="button" class="btn-link">
                   ASF
                 </button>
                 <ul>
                     <li>
                       <a href=https://apache.org/>Apache Homepage</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/licenses/>License</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/foundation/sponsorship.html>Sponsorship</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/security/>Security</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/foundation/thanks.html>Thanks</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/events/current-event>Events</a>
                     </li>
                     <li>
                       <a href=https://www.zhihu.com/column/c_1429578595417563136>Zhihu</a>
                     </li>
                 </ul>
               </div>
          </div>
            <div class="responsiveMenuIcon">
              <button type="button" id="menuBtn" class="btn-menu"><img src="../../_static/img/menu-icon.svg" alt="Menu Icon"></button>
            </div>

            <div class="tlcDropdown">
              <div class="dropdown">
                <button type="button" class="btn-link dropdown-toggle" data-toggle="dropdown" aria-haspopup="true" aria-expanded="false">
                  ASF
                </button>
                <div class="dropdown-menu dropdown-menu-right">
                  <ul>
                     <li>
                       <a href=https://apache.org/>Apache Homepage</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/licenses/>License</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/foundation/sponsorship.html>Sponsorship</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/security/>Security</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/foundation/thanks.html>Thanks</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/events/current-event>Events</a>
                     </li>
                     <li>
                       <a href=https://www.zhihu.com/column/c_1429578595417563136>Zhihu</a>
                     </li>
                  </ul>
                </div>
              </div>
          </div>
       </div>
    </div>
 </header>
 
    <nav data-toggle="wy-nav-shift" class="wy-nav-side fixed">
      <div class="wy-side-scroll">
        <div class="wy-side-nav-search" >
          

          
            <a href="../../index.html">
          

          
            
            <img src="../../_static/tvm-logo-small.png" class="logo" alt="Logo"/>
          
          </a>

          
            
            
                <div class="version">
                  0.8.dev1982
                </div>
            
          

          
<div role="search">
  <form id="rtd-search-form" class="wy-form" action="../../search.html" method="get">
    <input type="text" name="q" placeholder="Search docs" />
    <input type="hidden" name="check_keywords" value="yes" />
    <input type="hidden" name="area" value="default" />
  </form>
</div>

          
        </div>

        
        <div class="wy-menu wy-menu-vertical" data-spy="affix" role="navigation" aria-label="main navigation">
          
            
            
              
            
            
              <p class="caption" role="heading"><span class="caption-text">如何开始</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="../../install/index.html">安装 TVM</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../contribute/index.html">贡献者指南</a></li>
</ul>
<p class="caption" role="heading"><span class="caption-text">用户引导</span></p>
<ul class="current">
<li class="toctree-l1"><a class="reference internal" href="../../tutorial/index.html">User Tutorial</a></li>
<li class="toctree-l1 current"><a class="reference internal" href="../index.html">How To Guides</a><ul class="current">
<li class="toctree-l2"><a class="reference internal" href="../compile_models/index.html">编译深度学习模型</a></li>
<li class="toctree-l2"><a class="reference internal" href="../deploy/index.html">TVM 部署模型和集成</a></li>
<li class="toctree-l2"><a class="reference internal" href="../work_with_relay/index.html">Work With Relay</a></li>
<li class="toctree-l2"><a class="reference internal" href="../work_with_schedules/index.html">Work With Tensor Expression and Schedules</a></li>
<li class="toctree-l2 current"><a class="reference internal" href="index.html">优化张量算子</a><ul class="current">
<li class="toctree-l3"><a class="reference internal" href="opt_gemm.html">如何在CPU上优化GEMM（通用矩阵乘）</a></li>
<li class="toctree-l3 current"><a class="current reference internal" href="#">如何在GPU上优化卷积</a><ul>
<li class="toctree-l4"><a class="reference internal" href="#preparation-and-algorithm">准备和算法</a></li>
<li class="toctree-l4"><a class="reference internal" href="#memory-hierarchy">内存层次结构</a></li>
<li class="toctree-l4"><a class="reference internal" href="#blocking">分块</a></li>
<li class="toctree-l4"><a class="reference internal" href="#virtual-thread-split">虚拟线程拆分</a></li>
<li class="toctree-l4"><a class="reference internal" href="#cooperative-fetching">Cooperative Fetching</a></li>
<li class="toctree-l4"><a class="reference internal" href="#generate-cuda-kernel">Generate CUDA Kernel</a></li>
</ul>
</li>
<li class="toctree-l3"><a class="reference internal" href="opt_conv_tensorcore.html">How to optimize convolution using TensorCores</a></li>
</ul>
</li>
<li class="toctree-l2"><a class="reference internal" href="../tune_with_autotvm/index.html">Auto-Tune with Templates and AutoTVM</a></li>
<li class="toctree-l2"><a class="reference internal" href="../tune_with_autoscheduler/index.html">Use AutoScheduler for Template-Free Scheduling</a></li>
<li class="toctree-l2"><a class="reference internal" href="../work_with_microtvm/index.html">Work With microTVM</a></li>
<li class="toctree-l2"><a class="reference internal" href="../extend_tvm/index.html">Extend TVM</a></li>
<li class="toctree-l2"><a class="reference internal" href="../profile/index.html">Profile Models</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../errors.html">Handle TVM Errors</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../faq.html">常见提问</a></li>
</ul>
</li>
</ul>
<p class="caption" role="heading"><span class="caption-text">开发者引导</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="../../dev/tutorial/index.html">Developer Tutorial</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../dev/how_to/how_to.html">开发者指南</a></li>
</ul>
<p class="caption" role="heading"><span class="caption-text">架构指南</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="../../arch/index.html">Design and Architecture</a></li>
</ul>
<p class="caption" role="heading"><span class="caption-text">主题引导</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="../../topic/microtvm/index.html">microTVM：裸机使用TVM</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../topic/vta/index.html">VTA: Versatile Tensor Accelerator</a></li>
</ul>
<p class="caption" role="heading"><span class="caption-text">参考指南</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="../../reference/langref/index.html">语言参考</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../reference/api/python/index.html">Python API</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../reference/api/links.html">Other APIs</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../reference/publications.html">Publications</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../genindex.html">索引</a></li>
</ul>

            
          
        </div>
        
      </div>
    </nav>

    <section data-toggle="wy-nav-shift" class="wy-nav-content-wrap">
      
      <nav class="wy-nav-top" aria-label="top navigation" data-toggle="wy-nav-top">
        
            <div class="togglemenu">

            </div>
            <div class="nav-content">
              <!-- tvm -->
              Table of content
            </div>
        
      </nav>


      <div class="wy-nav-content">
        
        <div class="rst-content">
        

          




















<div role="navigation" aria-label="breadcrumbs navigation">

  <ul class="wy-breadcrumbs">
    
      <li><a href="../../index.html">Docs</a> <span class="br-arrow">></span></li>
        
          <li><a href="../index.html">How To Guides</a> <span class="br-arrow">></span></li>
        
          <li><a href="index.html">优化张量算子</a> <span class="br-arrow">></span></li>
        
      <li>如何在GPU上优化卷积</li>
    
    
      <li class="wy-breadcrumbs-aside">
        
            
            <a href="../../_sources/how_to/optimize_operators/opt_conv_cuda.rst.txt" rel="nofollow"> <img src="../../_static//img/source.svg" alt="viewsource"/></a>
          
        
      </li>
    
  </ul>

  
  <hr/>
</div>
          <div role="main" class="document" itemscope="itemscope" itemtype="http://schema.org/Article">
           <div itemprop="articleBody">
            
  <div class="sphx-glr-download-link-note admonition note">
<p class="admonition-title">注解</p>
<p>Click <a class="reference internal" href="#sphx-glr-download-how-to-optimize-operators-opt-conv-cuda-py"><span class="std std-ref">here</span></a> to download the full example code</p>
</div>
<div class="sphx-glr-example-title section" id="how-to-optimize-convolution-on-gpu">
<span id="opt-conv-gpu"></span><span id="sphx-glr-how-to-optimize-operators-opt-conv-cuda-py"></span><h1>如何在GPU上优化卷积<a class="headerlink" href="#how-to-optimize-convolution-on-gpu" title="永久链接至标题">¶</a></h1>
<p><strong>作者</strong>: <a class="reference external" href="https://homes.cs.washington.edu/~haichen/">Haichen Shen</a></p>
<p>在本教程中，我们将演示如何在 TVM 中编写高性能卷积实现。 我们以方形大小的输入张量和卷积核为例，并假设卷积的输入具有大的批量。 在这个例子中为了实现更好的数据局部性我们使用了一个不同的数据布局来存储数据。 缓冲区布局是 HWCN，代表高度、宽度、通道、批量。</p>
<div class="section" id="preparation-and-algorithm">
<h2>准备和算法<a class="headerlink" href="#preparation-and-algorithm" title="永久链接至标题">¶</a></h2>
<p>我们使用拥有固定尺寸的输入张量，其通道数为256，大小为14x14。批量大小为256。卷积核的大小为3x3，通道数有512个。我们将卷积层的步幅的大小和填充大小都设为1。以下代码定义了 TVM 中的卷积算法。</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="kn">import</span> <span class="nn">numpy</span> <span class="k">as</span> <span class="nn">np</span>
<span class="kn">import</span> <span class="nn">tvm</span>
<span class="kn">from</span> <span class="nn">tvm</span> <span class="k">import</span> <span class="n">te</span>

<span class="c1"># The sizes of inputs and filters</span>
<span class="n">batch</span> <span class="o">=</span> <span class="mi">256</span>
<span class="n">in_channel</span> <span class="o">=</span> <span class="mi">256</span>
<span class="n">out_channel</span> <span class="o">=</span> <span class="mi">512</span>
<span class="n">in_size</span> <span class="o">=</span> <span class="mi">14</span>
<span class="n">kernel</span> <span class="o">=</span> <span class="mi">3</span>
<span class="n">pad</span> <span class="o">=</span> <span class="mi">1</span>
<span class="n">stride</span> <span class="o">=</span> <span class="mi">1</span>

<span class="c1"># Algorithm</span>
<span class="n">A</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">placeholder</span><span class="p">((</span><span class="n">in_size</span><span class="p">,</span> <span class="n">in_size</span><span class="p">,</span> <span class="n">in_channel</span><span class="p">,</span> <span class="n">batch</span><span class="p">),</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;A&quot;</span><span class="p">)</span>
<span class="n">W</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">placeholder</span><span class="p">((</span><span class="n">kernel</span><span class="p">,</span> <span class="n">kernel</span><span class="p">,</span> <span class="n">in_channel</span><span class="p">,</span> <span class="n">out_channel</span><span class="p">),</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;W&quot;</span><span class="p">)</span>
<span class="n">out_size</span> <span class="o">=</span> <span class="p">(</span><span class="n">in_size</span> <span class="o">-</span> <span class="n">kernel</span> <span class="o">+</span> <span class="mi">2</span> <span class="o">*</span> <span class="n">pad</span><span class="p">)</span> <span class="o">//</span> <span class="n">stride</span> <span class="o">+</span> <span class="mi">1</span>
<span class="c1"># Pad input</span>
<span class="n">Apad</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">compute</span><span class="p">(</span>
    <span class="p">(</span><span class="n">in_size</span> <span class="o">+</span> <span class="mi">2</span> <span class="o">*</span> <span class="n">pad</span><span class="p">,</span> <span class="n">in_size</span> <span class="o">+</span> <span class="mi">2</span> <span class="o">*</span> <span class="n">pad</span><span class="p">,</span> <span class="n">in_channel</span><span class="p">,</span> <span class="n">batch</span><span class="p">),</span>
    <span class="k">lambda</span> <span class="n">yy</span><span class="p">,</span> <span class="n">xx</span><span class="p">,</span> <span class="n">cc</span><span class="p">,</span> <span class="n">nn</span><span class="p">:</span> <span class="n">tvm</span><span class="o">.</span><span class="n">tir</span><span class="o">.</span><span class="n">if_then_else</span><span class="p">(</span>
        <span class="n">tvm</span><span class="o">.</span><span class="n">tir</span><span class="o">.</span><span class="n">all</span><span class="p">(</span><span class="n">yy</span> <span class="o">&gt;=</span> <span class="n">pad</span><span class="p">,</span> <span class="n">yy</span> <span class="o">-</span> <span class="n">pad</span> <span class="o">&lt;</span> <span class="n">in_size</span><span class="p">,</span> <span class="n">xx</span> <span class="o">&gt;=</span> <span class="n">pad</span><span class="p">,</span> <span class="n">xx</span> <span class="o">-</span> <span class="n">pad</span> <span class="o">&lt;</span> <span class="n">in_size</span><span class="p">),</span>
        <span class="n">A</span><span class="p">[</span><span class="n">yy</span> <span class="o">-</span> <span class="n">pad</span><span class="p">,</span> <span class="n">xx</span> <span class="o">-</span> <span class="n">pad</span><span class="p">,</span> <span class="n">cc</span><span class="p">,</span> <span class="n">nn</span><span class="p">],</span>
        <span class="n">tvm</span><span class="o">.</span><span class="n">tir</span><span class="o">.</span><span class="n">const</span><span class="p">(</span><span class="mf">0.0</span><span class="p">,</span> <span class="s2">&quot;float32&quot;</span><span class="p">),</span>
    <span class="p">),</span>
    <span class="n">name</span><span class="o">=</span><span class="s2">&quot;Apad&quot;</span><span class="p">,</span>
<span class="p">)</span>
<span class="c1"># Create reduction variables</span>
<span class="n">rc</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">reduce_axis</span><span class="p">((</span><span class="mi">0</span><span class="p">,</span> <span class="n">in_channel</span><span class="p">),</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;rc&quot;</span><span class="p">)</span>
<span class="n">ry</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">reduce_axis</span><span class="p">((</span><span class="mi">0</span><span class="p">,</span> <span class="n">kernel</span><span class="p">),</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;ry&quot;</span><span class="p">)</span>
<span class="n">rx</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">reduce_axis</span><span class="p">((</span><span class="mi">0</span><span class="p">,</span> <span class="n">kernel</span><span class="p">),</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;rx&quot;</span><span class="p">)</span>
<span class="c1"># Compute the convolution</span>
<span class="n">B</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">compute</span><span class="p">(</span>
    <span class="p">(</span><span class="n">out_size</span><span class="p">,</span> <span class="n">out_size</span><span class="p">,</span> <span class="n">out_channel</span><span class="p">,</span> <span class="n">batch</span><span class="p">),</span>
    <span class="k">lambda</span> <span class="n">yy</span><span class="p">,</span> <span class="n">xx</span><span class="p">,</span> <span class="n">ff</span><span class="p">,</span> <span class="n">nn</span><span class="p">:</span> <span class="n">te</span><span class="o">.</span><span class="n">sum</span><span class="p">(</span>
        <span class="n">Apad</span><span class="p">[</span><span class="n">yy</span> <span class="o">*</span> <span class="n">stride</span> <span class="o">+</span> <span class="n">ry</span><span class="p">,</span> <span class="n">xx</span> <span class="o">*</span> <span class="n">stride</span> <span class="o">+</span> <span class="n">rx</span><span class="p">,</span> <span class="n">rc</span><span class="p">,</span> <span class="n">nn</span><span class="p">]</span> <span class="o">*</span> <span class="n">W</span><span class="p">[</span><span class="n">ry</span><span class="p">,</span> <span class="n">rx</span><span class="p">,</span> <span class="n">rc</span><span class="p">,</span> <span class="n">ff</span><span class="p">],</span> <span class="n">axis</span><span class="o">=</span><span class="p">[</span><span class="n">ry</span><span class="p">,</span> <span class="n">rx</span><span class="p">,</span> <span class="n">rc</span><span class="p">]</span>
    <span class="p">),</span>
    <span class="n">name</span><span class="o">=</span><span class="s2">&quot;B&quot;</span><span class="p">,</span>
<span class="p">)</span>
</pre></div>
</div>
</div>
<div class="section" id="memory-hierarchy">
<h2>内存层次结构<a class="headerlink" href="#memory-hierarchy" title="永久链接至标题">¶</a></h2>
<p>我们首先指定缓冲区的内存层次结构。 下图显示了 GPU 内存层次结构。 与 CPU 内存层次结构的一个重要区别是 GPU 提供了一个称为共享内存的缓存缓冲区，由程序员来管理。
因此，要实现高性能的GPU内核，最大化共享内存中的数据重用是至关重要的。</p>
<a class="reference internal image-reference" href="https://github.com/dmlc/web-data/raw/main/tvm/tutorial/gpu_memory_hierarchy.png"><img alt="https://github.com/dmlc/web-data/raw/main/tvm/tutorial/gpu_memory_hierarchy.png" class="align-center" src="https://github.com/dmlc/web-data/raw/main/tvm/tutorial/gpu_memory_hierarchy.png" style="width: 271px; height: 319px;" /></a>
<p>在这个例子中，我们将 Apad 和 W 加载到缓冲区 AA 和 WW 中，它们存储在共享内存中。 这些缓冲区稍后将由同一线程块内的所有线程共享以计算卷积。 然后每个线程从共享缓冲区加载它自己的部分到它们的本地寄存器 AL 和 WL中。 BL 是输出 B 的本地缓存，也存储在线程本地寄存器中。</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="c1"># Designate the memory hierarchy</span>
<span class="n">s</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">create_schedule</span><span class="p">(</span><span class="n">B</span><span class="o">.</span><span class="n">op</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">Apad</span><span class="p">]</span><span class="o">.</span><span class="n">compute_inline</span><span class="p">()</span>  <span class="c1"># compute Apad inline</span>
<span class="n">AA</span> <span class="o">=</span> <span class="n">s</span><span class="o">.</span><span class="n">cache_read</span><span class="p">(</span><span class="n">Apad</span><span class="p">,</span> <span class="s2">&quot;shared&quot;</span><span class="p">,</span> <span class="p">[</span><span class="n">B</span><span class="p">])</span>
<span class="n">WW</span> <span class="o">=</span> <span class="n">s</span><span class="o">.</span><span class="n">cache_read</span><span class="p">(</span><span class="n">W</span><span class="p">,</span> <span class="s2">&quot;shared&quot;</span><span class="p">,</span> <span class="p">[</span><span class="n">B</span><span class="p">])</span>
<span class="n">AL</span> <span class="o">=</span> <span class="n">s</span><span class="o">.</span><span class="n">cache_read</span><span class="p">(</span><span class="n">AA</span><span class="p">,</span> <span class="s2">&quot;local&quot;</span><span class="p">,</span> <span class="p">[</span><span class="n">B</span><span class="p">])</span>
<span class="n">WL</span> <span class="o">=</span> <span class="n">s</span><span class="o">.</span><span class="n">cache_read</span><span class="p">(</span><span class="n">WW</span><span class="p">,</span> <span class="s2">&quot;local&quot;</span><span class="p">,</span> <span class="p">[</span><span class="n">B</span><span class="p">])</span>
<span class="n">BL</span> <span class="o">=</span> <span class="n">s</span><span class="o">.</span><span class="n">cache_write</span><span class="p">(</span><span class="n">B</span><span class="p">,</span> <span class="s2">&quot;local&quot;</span><span class="p">)</span>
</pre></div>
</div>
</div>
<div class="section" id="blocking">
<h2>分块<a class="headerlink" href="#blocking" title="永久链接至标题">¶</a></h2>
<p>以下代码将矩阵乘法的整个工作负载拆分为线程块和单个线程。 我们在矩阵乘法中遵循分块方案。 如下图所示，对于某个批量和通道数来说，给定一个像素坐标（y，x），一个线程块负责计算一个block_factor x block_factor（64 x 64）的区域。 由于共享内存空间的限制，我们每次只从 Apad 和 B 加载 step x block_factor (8 x 64) 数据到共享内存中的缓冲区。</p>
<a class="reference internal image-reference" href="https://github.com/dmlc/web-data/raw/main/tvm/tutorial/conv_gpu_blocking.png"><img alt="https://github.com/dmlc/web-data/raw/main/tvm/tutorial/conv_gpu_blocking.png" class="align-center" src="https://github.com/dmlc/web-data/raw/main/tvm/tutorial/conv_gpu_blocking.png" style="width: 317px; height: 308px;" /></a>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="c1"># tile consts</span>
<span class="n">tile</span> <span class="o">=</span> <span class="mi">8</span>
<span class="n">num_thread</span> <span class="o">=</span> <span class="mi">8</span>
<span class="n">block_factor</span> <span class="o">=</span> <span class="n">tile</span> <span class="o">*</span> <span class="n">num_thread</span>
<span class="n">step</span> <span class="o">=</span> <span class="mi">8</span>
<span class="n">vthread</span> <span class="o">=</span> <span class="mi">2</span>

<span class="c1"># Get the GPU thread indices</span>
<span class="n">block_x</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">thread_axis</span><span class="p">(</span><span class="s2">&quot;blockIdx.x&quot;</span><span class="p">)</span>
<span class="n">block_y</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">thread_axis</span><span class="p">(</span><span class="s2">&quot;blockIdx.y&quot;</span><span class="p">)</span>
<span class="n">block_z</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">thread_axis</span><span class="p">(</span><span class="s2">&quot;blockIdx.z&quot;</span><span class="p">)</span>
<span class="n">thread_x</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">thread_axis</span><span class="p">((</span><span class="mi">0</span><span class="p">,</span> <span class="n">num_thread</span><span class="p">),</span> <span class="s2">&quot;threadIdx.x&quot;</span><span class="p">)</span>
<span class="n">thread_y</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">thread_axis</span><span class="p">((</span><span class="mi">0</span><span class="p">,</span> <span class="n">num_thread</span><span class="p">),</span> <span class="s2">&quot;threadIdx.y&quot;</span><span class="p">)</span>
<span class="n">thread_xz</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">thread_axis</span><span class="p">((</span><span class="mi">0</span><span class="p">,</span> <span class="n">vthread</span><span class="p">),</span> <span class="s2">&quot;vthread&quot;</span><span class="p">,</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;vx&quot;</span><span class="p">)</span>
<span class="n">thread_yz</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">thread_axis</span><span class="p">((</span><span class="mi">0</span><span class="p">,</span> <span class="n">vthread</span><span class="p">),</span> <span class="s2">&quot;vthread&quot;</span><span class="p">,</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;vy&quot;</span><span class="p">)</span>

<span class="c1"># Split the workloads</span>
<span class="n">hi</span><span class="p">,</span> <span class="n">wi</span><span class="p">,</span> <span class="n">fi</span><span class="p">,</span> <span class="n">ni</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">B</span><span class="p">]</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">axis</span>
<span class="n">bz</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">B</span><span class="p">]</span><span class="o">.</span><span class="n">fuse</span><span class="p">(</span><span class="n">hi</span><span class="p">,</span> <span class="n">wi</span><span class="p">)</span>
<span class="n">by</span><span class="p">,</span> <span class="n">fi</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">B</span><span class="p">]</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="n">fi</span><span class="p">,</span> <span class="n">factor</span><span class="o">=</span><span class="n">block_factor</span><span class="p">)</span>
<span class="n">bx</span><span class="p">,</span> <span class="n">ni</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">B</span><span class="p">]</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="n">ni</span><span class="p">,</span> <span class="n">factor</span><span class="o">=</span><span class="n">block_factor</span><span class="p">)</span>

<span class="c1"># Bind the iteration variables to GPU thread indices</span>
<span class="n">s</span><span class="p">[</span><span class="n">B</span><span class="p">]</span><span class="o">.</span><span class="n">bind</span><span class="p">(</span><span class="n">bz</span><span class="p">,</span> <span class="n">block_z</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">B</span><span class="p">]</span><span class="o">.</span><span class="n">bind</span><span class="p">(</span><span class="n">by</span><span class="p">,</span> <span class="n">block_y</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">B</span><span class="p">]</span><span class="o">.</span><span class="n">bind</span><span class="p">(</span><span class="n">bx</span><span class="p">,</span> <span class="n">block_x</span><span class="p">)</span>
</pre></div>
</div>
</div>
<div class="section" id="virtual-thread-split">
<h2>虚拟线程拆分<a class="headerlink" href="#virtual-thread-split" title="永久链接至标题">¶</a></h2>
<p>We further split the workload from a thread block to individual threads. To
avoid <em>memory bank conflict</em>, we use virtual thread to split the area into 4
parts, and then tile into 8x8 grids. Therefore, shown in the figure below,
each thread computes 4 strided grids, where size of each grid is 4 x 4.</p>
<a class="reference internal image-reference" href="https://github.com/dmlc/web-data/raw/main/tvm/tutorial/conv_gpu_vthread.png"><img alt="https://github.com/dmlc/web-data/raw/main/tvm/tutorial/conv_gpu_vthread.png" class="align-center" src="https://github.com/dmlc/web-data/raw/main/tvm/tutorial/conv_gpu_vthread.png" style="width: 268px; height: 188px;" /></a>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="n">tyz</span><span class="p">,</span> <span class="n">fi</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">B</span><span class="p">]</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="n">fi</span><span class="p">,</span> <span class="n">nparts</span><span class="o">=</span><span class="n">vthread</span><span class="p">)</span>  <span class="c1"># virtual thread split</span>
<span class="n">txz</span><span class="p">,</span> <span class="n">ni</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">B</span><span class="p">]</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="n">ni</span><span class="p">,</span> <span class="n">nparts</span><span class="o">=</span><span class="n">vthread</span><span class="p">)</span>  <span class="c1"># virtual thread split</span>
<span class="n">ty</span><span class="p">,</span> <span class="n">fi</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">B</span><span class="p">]</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="n">fi</span><span class="p">,</span> <span class="n">nparts</span><span class="o">=</span><span class="n">num_thread</span><span class="p">)</span>
<span class="n">tx</span><span class="p">,</span> <span class="n">ni</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">B</span><span class="p">]</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="n">ni</span><span class="p">,</span> <span class="n">nparts</span><span class="o">=</span><span class="n">num_thread</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">B</span><span class="p">]</span><span class="o">.</span><span class="n">reorder</span><span class="p">(</span><span class="n">bz</span><span class="p">,</span> <span class="n">by</span><span class="p">,</span> <span class="n">bx</span><span class="p">,</span> <span class="n">tyz</span><span class="p">,</span> <span class="n">txz</span><span class="p">,</span> <span class="n">ty</span><span class="p">,</span> <span class="n">tx</span><span class="p">,</span> <span class="n">fi</span><span class="p">,</span> <span class="n">ni</span><span class="p">)</span>

<span class="n">s</span><span class="p">[</span><span class="n">B</span><span class="p">]</span><span class="o">.</span><span class="n">bind</span><span class="p">(</span><span class="n">tyz</span><span class="p">,</span> <span class="n">thread_yz</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">B</span><span class="p">]</span><span class="o">.</span><span class="n">bind</span><span class="p">(</span><span class="n">txz</span><span class="p">,</span> <span class="n">thread_xz</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">B</span><span class="p">]</span><span class="o">.</span><span class="n">bind</span><span class="p">(</span><span class="n">ty</span><span class="p">,</span> <span class="n">thread_y</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">B</span><span class="p">]</span><span class="o">.</span><span class="n">bind</span><span class="p">(</span><span class="n">tx</span><span class="p">,</span> <span class="n">thread_x</span><span class="p">)</span>
</pre></div>
</div>
</div>
<div class="section" id="cooperative-fetching">
<h2>Cooperative Fetching<a class="headerlink" href="#cooperative-fetching" title="永久链接至标题">¶</a></h2>
<p>As mentioned before, each time step we need to transfer step x block_factor
data from GPU global memory to shared memory. In order to reduce the memory
transfer per thread, the following code lets threads in the same thread block
coopertively fetch dependent data from global memory.</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="c1"># Schedule BL local write</span>
<span class="n">s</span><span class="p">[</span><span class="n">BL</span><span class="p">]</span><span class="o">.</span><span class="n">compute_at</span><span class="p">(</span><span class="n">s</span><span class="p">[</span><span class="n">B</span><span class="p">],</span> <span class="n">tx</span><span class="p">)</span>
<span class="n">yi</span><span class="p">,</span> <span class="n">xi</span><span class="p">,</span> <span class="n">fi</span><span class="p">,</span> <span class="n">ni</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">BL</span><span class="p">]</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">axis</span>
<span class="n">ry</span><span class="p">,</span> <span class="n">rx</span><span class="p">,</span> <span class="n">rc</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">BL</span><span class="p">]</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">reduce_axis</span>
<span class="n">rco</span><span class="p">,</span> <span class="n">rci</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">BL</span><span class="p">]</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="n">rc</span><span class="p">,</span> <span class="n">factor</span><span class="o">=</span><span class="n">step</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">BL</span><span class="p">]</span><span class="o">.</span><span class="n">reorder</span><span class="p">(</span><span class="n">rco</span><span class="p">,</span> <span class="n">ry</span><span class="p">,</span> <span class="n">rx</span><span class="p">,</span> <span class="n">rci</span><span class="p">,</span> <span class="n">fi</span><span class="p">,</span> <span class="n">ni</span><span class="p">)</span>

<span class="c1"># Attach computation to iteration variables</span>
<span class="n">s</span><span class="p">[</span><span class="n">AA</span><span class="p">]</span><span class="o">.</span><span class="n">compute_at</span><span class="p">(</span><span class="n">s</span><span class="p">[</span><span class="n">BL</span><span class="p">],</span> <span class="n">rx</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">WW</span><span class="p">]</span><span class="o">.</span><span class="n">compute_at</span><span class="p">(</span><span class="n">s</span><span class="p">[</span><span class="n">BL</span><span class="p">],</span> <span class="n">rx</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">AL</span><span class="p">]</span><span class="o">.</span><span class="n">compute_at</span><span class="p">(</span><span class="n">s</span><span class="p">[</span><span class="n">BL</span><span class="p">],</span> <span class="n">rci</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">WL</span><span class="p">]</span><span class="o">.</span><span class="n">compute_at</span><span class="p">(</span><span class="n">s</span><span class="p">[</span><span class="n">BL</span><span class="p">],</span> <span class="n">rci</span><span class="p">)</span>

<span class="c1"># Schedule for A&#39;s shared memory load</span>
<span class="n">yi</span><span class="p">,</span> <span class="n">xi</span><span class="p">,</span> <span class="n">ci</span><span class="p">,</span> <span class="n">ni</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">AA</span><span class="p">]</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">axis</span>
<span class="n">ty</span><span class="p">,</span> <span class="n">ci</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">AA</span><span class="p">]</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="n">ci</span><span class="p">,</span> <span class="n">nparts</span><span class="o">=</span><span class="n">num_thread</span><span class="p">)</span>
<span class="n">tx</span><span class="p">,</span> <span class="n">ni</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">AA</span><span class="p">]</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="n">ni</span><span class="p">,</span> <span class="n">nparts</span><span class="o">=</span><span class="n">num_thread</span><span class="p">)</span>
<span class="n">_</span><span class="p">,</span> <span class="n">ni</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">AA</span><span class="p">]</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="n">ni</span><span class="p">,</span> <span class="n">factor</span><span class="o">=</span><span class="mi">4</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">AA</span><span class="p">]</span><span class="o">.</span><span class="n">reorder</span><span class="p">(</span><span class="n">ty</span><span class="p">,</span> <span class="n">tx</span><span class="p">,</span> <span class="n">yi</span><span class="p">,</span> <span class="n">xi</span><span class="p">,</span> <span class="n">ci</span><span class="p">,</span> <span class="n">ni</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">AA</span><span class="p">]</span><span class="o">.</span><span class="n">bind</span><span class="p">(</span><span class="n">ty</span><span class="p">,</span> <span class="n">thread_y</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">AA</span><span class="p">]</span><span class="o">.</span><span class="n">bind</span><span class="p">(</span><span class="n">tx</span><span class="p">,</span> <span class="n">thread_x</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">AA</span><span class="p">]</span><span class="o">.</span><span class="n">vectorize</span><span class="p">(</span><span class="n">ni</span><span class="p">)</span>  <span class="c1"># vectorize memory load</span>

<span class="c1"># Schedule for W&#39;s shared memory load</span>
<span class="n">yi</span><span class="p">,</span> <span class="n">xi</span><span class="p">,</span> <span class="n">ci</span><span class="p">,</span> <span class="n">fi</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">WW</span><span class="p">]</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">axis</span>
<span class="n">ty</span><span class="p">,</span> <span class="n">ci</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">WW</span><span class="p">]</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="n">ci</span><span class="p">,</span> <span class="n">nparts</span><span class="o">=</span><span class="n">num_thread</span><span class="p">)</span>
<span class="n">tx</span><span class="p">,</span> <span class="n">fi</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">WW</span><span class="p">]</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="n">fi</span><span class="p">,</span> <span class="n">nparts</span><span class="o">=</span><span class="n">num_thread</span><span class="p">)</span>
<span class="n">_</span><span class="p">,</span> <span class="n">fi</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">WW</span><span class="p">]</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="n">fi</span><span class="p">,</span> <span class="n">factor</span><span class="o">=</span><span class="mi">4</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">WW</span><span class="p">]</span><span class="o">.</span><span class="n">reorder</span><span class="p">(</span><span class="n">ty</span><span class="p">,</span> <span class="n">tx</span><span class="p">,</span> <span class="n">yi</span><span class="p">,</span> <span class="n">xi</span><span class="p">,</span> <span class="n">ci</span><span class="p">,</span> <span class="n">fi</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">WW</span><span class="p">]</span><span class="o">.</span><span class="n">bind</span><span class="p">(</span><span class="n">ty</span><span class="p">,</span> <span class="n">thread_y</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">WW</span><span class="p">]</span><span class="o">.</span><span class="n">bind</span><span class="p">(</span><span class="n">tx</span><span class="p">,</span> <span class="n">thread_x</span><span class="p">)</span>
<span class="n">s</span><span class="p">[</span><span class="n">WW</span><span class="p">]</span><span class="o">.</span><span class="n">vectorize</span><span class="p">(</span><span class="n">fi</span><span class="p">)</span>  <span class="c1"># vectorize memory load</span>
</pre></div>
</div>
</div>
<div class="section" id="generate-cuda-kernel">
<h2>Generate CUDA Kernel<a class="headerlink" href="#generate-cuda-kernel" title="永久链接至标题">¶</a></h2>
<p>Finally we use TVM to generate and compile the CUDA kernel, and evaluate the
latency of convolution.</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="n">func</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">build</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="p">[</span><span class="n">A</span><span class="p">,</span> <span class="n">W</span><span class="p">,</span> <span class="n">B</span><span class="p">],</span> <span class="s2">&quot;cuda&quot;</span><span class="p">)</span>
<span class="n">dev</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">cuda</span><span class="p">(</span><span class="mi">0</span><span class="p">)</span>
<span class="n">a_np</span> <span class="o">=</span> <span class="n">np</span><span class="o">.</span><span class="n">random</span><span class="o">.</span><span class="n">uniform</span><span class="p">(</span><span class="n">size</span><span class="o">=</span><span class="p">(</span><span class="n">in_size</span><span class="p">,</span> <span class="n">in_size</span><span class="p">,</span> <span class="n">in_channel</span><span class="p">,</span> <span class="n">batch</span><span class="p">))</span><span class="o">.</span><span class="n">astype</span><span class="p">(</span><span class="n">A</span><span class="o">.</span><span class="n">dtype</span><span class="p">)</span>
<span class="n">w_np</span> <span class="o">=</span> <span class="n">np</span><span class="o">.</span><span class="n">random</span><span class="o">.</span><span class="n">uniform</span><span class="p">(</span><span class="n">size</span><span class="o">=</span><span class="p">(</span><span class="n">kernel</span><span class="p">,</span> <span class="n">kernel</span><span class="p">,</span> <span class="n">in_channel</span><span class="p">,</span> <span class="n">out_channel</span><span class="p">))</span><span class="o">.</span><span class="n">astype</span><span class="p">(</span><span class="n">W</span><span class="o">.</span><span class="n">dtype</span><span class="p">)</span>
<span class="n">a</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span><span class="p">(</span><span class="n">a_np</span><span class="p">,</span> <span class="n">dev</span><span class="p">)</span>
<span class="n">w</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span><span class="p">(</span><span class="n">w_np</span><span class="p">,</span> <span class="n">dev</span><span class="p">)</span>
<span class="n">b</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">zeros</span><span class="p">((</span><span class="n">out_size</span><span class="p">,</span> <span class="n">out_size</span><span class="p">,</span> <span class="n">out_channel</span><span class="p">,</span> <span class="n">batch</span><span class="p">),</span> <span class="n">dtype</span><span class="o">=</span><span class="n">B</span><span class="o">.</span><span class="n">dtype</span><span class="p">),</span> <span class="n">dev</span><span class="p">)</span>
<span class="n">func</span><span class="p">(</span><span class="n">a</span><span class="p">,</span> <span class="n">w</span><span class="p">,</span> <span class="n">b</span><span class="p">)</span>
<span class="n">evaluator</span> <span class="o">=</span> <span class="n">func</span><span class="o">.</span><span class="n">time_evaluator</span><span class="p">(</span><span class="n">func</span><span class="o">.</span><span class="n">entry_name</span><span class="p">,</span> <span class="n">dev</span><span class="p">,</span> <span class="n">number</span><span class="o">=</span><span class="mi">1</span><span class="p">)</span>
<span class="nb">print</span><span class="p">(</span><span class="s2">&quot;Convolution: </span><span class="si">%f</span><span class="s2"> ms&quot;</span> <span class="o">%</span> <span class="p">(</span><span class="n">evaluator</span><span class="p">(</span><span class="n">a</span><span class="p">,</span> <span class="n">w</span><span class="p">,</span> <span class="n">b</span><span class="p">)</span><span class="o">.</span><span class="n">mean</span> <span class="o">*</span> <span class="mf">1e3</span><span class="p">))</span>
</pre></div>
</div>
<p class="sphx-glr-script-out">输出:</p>
<div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>Convolution: 14.042834 ms
</pre></div>
</div>
<div class="sphx-glr-footer class sphx-glr-footer-example docutils container" id="sphx-glr-download-how-to-optimize-operators-opt-conv-cuda-py">
<div class="sphx-glr-download docutils container">
<p><a class="reference download internal" download="" href="../../_downloads/3c5c85c3954f3110f16ca084e286f03a/opt_conv_cuda.py"><code class="xref download docutils literal notranslate"><span class="pre">下载</span> <span class="pre">Python</span> <span class="pre">源代码:</span> <span class="pre">opt_conv_cuda.py</span></code></a></p>
</div>
<div class="sphx-glr-download docutils container">
<p><a class="reference download internal" download="" href="../../_downloads/854257a66df713b1f3f82eb3577f95e3/opt_conv_cuda.ipynb"><code class="xref download docutils literal notranslate"><span class="pre">下载</span> <span class="pre">Jupyter</span> <span class="pre">notebook:</span> <span class="pre">opt_conv_cuda.ipynb</span></code></a></p>
</div>
</div>
<p class="sphx-glr-signature"><a class="reference external" href="https://sphinx-gallery.github.io">Gallery generated by Sphinx-Gallery</a></p>
</div>
</div>


           </div>
           
          </div>
          

<footer>

    <div class="rst-footer-buttons" role="navigation" aria-label="footer navigation">
      
        <a href="opt_conv_tensorcore.html" class="btn btn-neutral float-right" title="How to optimize convolution using TensorCores" accesskey="n" rel="next">下一个 <span class="fa fa-arrow-circle-right"></span></a>
      
      
        <a href="opt_gemm.html" class="btn btn-neutral float-left" title="如何在CPU上优化GEMM（通用矩阵乘）" accesskey="p" rel="prev"><span class="fa fa-arrow-circle-left"></span> 上一个</a>
      
    </div>

<div id="button" class="backtop"><img src="../../_static//img/right.svg" alt="backtop"/> </div>
<section class="footerSec">
    <div class="footerHeader">
      <ul class="d-flex align-md-items-center justify-content-between flex-column flex-md-row">
        <li class="copywrite d-flex align-items-center">
          <h5 id="copy-right-info">© 2020 Apache Software Foundation | All right reserved</h5>
        </li>
      </ul>

    </div>

    <ul>
      <li class="footernote">Copyright © 2020 The Apache Software Foundation. Apache TVM, Apache, the Apache feather, and the Apache TVM project logo are either trademarks or registered trademarks of the Apache Software Foundation.</li>
    </ul>

</section>
</footer>
        </div>
      </div>

    </section>

  </div>
  

    <script src="https://cdnjs.cloudflare.com/ajax/libs/popper.js/1.12.9/umd/popper.min.js" integrity="sha384-ApNbgh9B+Y1QKtv3Rn7W3mgPxhU9K/ScQsAP7hUibX39j7fakFPskvXusvfa0b4Q" crossorigin="anonymous"></script>
    <script src="https://maxcdn.bootstrapcdn.com/bootstrap/4.0.0/js/bootstrap.min.js" integrity="sha384-JZR6Spejh4U02d8jOt6vLEHfe/JQGiRRSQQxSfFWpi1MquVdAyjUar5+76PVCmYl" crossorigin="anonymous"></script>

  </body>
  <script type="text/javascript">
      jQuery(function () {
          SphinxRtdTheme.Navigation.enable(true);
      });
  </script>

  
  
    
    <!-- Theme Analytics -->
    <script>
    (function(i,s,o,g,r,a,m){i['GoogleAnalyticsObject']=r;i[r]=i[r]||function(){
      (i[r].q=i[r].q||[]).push(arguments)},i[r].l=1*new Date();a=s.createElement(o),
      m=s.getElementsByTagName(o)[0];a.async=1;a.src=g;m.parentNode.insertBefore(a,m)
    })(window,document,'script','https://www.google-analytics.com/analytics.js','ga');

    ga('create', 'UA-75982049-2', 'auto');
    ga('send', 'pageview');
    </script>

    
   

</body>
</html>